#ifndef	__cuMath_cuh__
#define	__cuMath_cuh__

#include <cuda.h>
#include <vector_functions.h>
#include <math_functions.h>

typedef	unsigned	int		uint;

#ifndef	FALSE
#	define	FALSE	0
#endif	// FALSE

#ifndef	TRUE
#	define	TRUE	1
#endif	// TRUE

#define	CU_EPSILON		1e-6f

__device__ __inline__ unsigned int RGBAFloatToUint(float4 rgba)
{
	return (unsigned int(rgba.w * 255.0f) << 24) |
		   (unsigned int(rgba.z * 255.0f) << 16) |
		   (unsigned int(rgba.y * 255.0f) <<  8) |
			unsigned int(rgba.x * 255.0f);
}
__device__ __inline__ float clamp(float a, float minimum, float maximum)
{
	return ((a < minimum) ? minimum : ((a > maximum) ? maximum : a));
}
__device__ __inline__ float signf(const float val)
{
	return (val < 0.0f) ? -1.0f : 1.0f;
}
//__inline__ __device__ float  saturate(float a)				{ return clamp(x, 0.0f, 1.0f); }
__device__ __inline__ float2 float2_sub(float2 a, float2 b)		{ return make_float2(a.x - b.x, a.y - b.y); }

__device__ __inline__ float3 float3_sub(float3 a, float3 b)		{ return make_float3(a.x - b.x, a.y - b.y, a.z - b.z); }
__device__ __inline__ float3 float3_sub(float3 a, float  s)		{ return make_float3(a.x -   s, a.y -   s, a.z -   s); }
__device__ __inline__ float3 float3_add(float3 a, float3 b)		{ return make_float3(a.x + b.x, a.y + b.y, a.z + b.z); }
__device__ __inline__ float3 float3_mul(float3 a, float3 b)		{ return make_float3(a.x * b.x, a.y * b.y, a.z * b.z); }
__device__ __inline__ float3 float3_mul(float3 a, float  s)		{ return make_float3(a.x *   s, a.y *   s, a.z *   s); }
__device__ __inline__ float3 float3_div(float3 a, float3 b)		{ return make_float3(a.x / b.x, a.y / b.y, a.z / b.z); }
__device__ __inline__ float3 float3_div(float3 a, float  s)		{ float invS = 1.0f / s; return make_float3(a.x * invS, a.y * invS, a.z * invS); }

__device__ __inline__ float4 make_float4(float3 xyz, float w)	{ return make_float4(xyz.x, xyz.y, xyz.z, w); }
__device__ __inline__ float4 float4_add(float4 a, float4 b)		{ return make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); }
__device__ __inline__ float4 float4_mul(float4 a, float4 b)		{ return make_float4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); }
__device__ __inline__ float4 float4_mul(float4 a, float s)		{ return make_float4(a.x *   s, a.y *   s, a.z *   s, a.w *   s); }

__device__ __inline__ float3 saturate(float3 a)					{ return make_float3(saturate(a.x), saturate(a.y), saturate(a.z)); }
__device__ __inline__ float4 saturate(float4 a)					{ return make_float4(saturate(a.x), saturate(a.y), saturate(a.z), saturate(a.w)); }
__device__ __inline__ float  dot(float3 a, float3 b)			{ return a.x * b.x + a.y * b.y + a.z * b.z; }
__device__ __inline__ float  dot(float4 a, float4 b)			{ return a.x * b.x + a.y * b.y + a.z * b.z + a.w * b.w; }
__device__ __inline__ float  length(float3 v)					{ return sqrtf(dot(v, v)); }
__device__ __inline__ float3 normalize(float3 v)				{ float inv_len = rsqrtf(dot(v, v)); return float3_mul(v, inv_len); }
__device__ __inline__ float3 reflect(float3 dir, float3 normal)
{
	float3 temp = float3_mul(float3_mul(normal, dot(dir, normal)), 2.0f);
	return normalize(float3_sub(dir, temp));
}

__device__ __inline__ float3 transform_by_normal(const float3& normal, const float3& v)
{
	float sinV = -v.y;
	float cosV = sqrtf(1.0f - sinV * sinV);

	float cosU, sinU;

	if (cosV == 0.0f)
	{
		sinU = 0.0f;
		cosU = 1.0f;
	}
	else
	{
		float invCosV = 1.0f / cosV;
		sinU = v.x * invCosV;
		cosU = v.z * invCosV;
	}

	return make_float3((normal.y * sinV + normal.z * cosV) * sinU + normal.x * cosU,
						normal.y * cosV - normal.z * sinV,
					   (normal.y * sinV + normal.z * cosV) * cosU - normal.x * sinU);
}


struct Ray
{
	__device__ Ray()
		: weight(1.0f)
	{
	}
	__device__ Ray(float3 o, float3 d)
		: origin(o)
		, weight(1.0f)
	{
		SetDir(d);
	}
	__device__ Ray(const Ray& other)
	{
		(*this) = other;
	}

	__device__ __inline__ void SetDir(const float3& _dir)
	{
		dir = normalize(_dir);
		dir_inv = make_float3(1.0f / dir.x, 1.0f / dir.y, 1.0f / dir.z);
		posneg[0] = (dir.x > 0 ? 0 : 1);
		posneg[1] = (dir.y > 0 ? 0 : 1);
		posneg[2] = (dir.z > 0 ? 0 : 1);
	}

	__device__ __inline__ const Ray& operator = (const Ray& other)
	{
		origin = other.origin;
		dir = other.dir;
		dir_inv = other.dir_inv;
		weight = other.weight;
		posneg[0] = other.posneg[0];
		posneg[1] = other.posneg[1];
		posneg[2] = other.posneg[2];

		return *this;
	}

	float3		origin;
	float3		dir;
	float3		dir_inv;
	float		weight;
	int			posneg[3];
};

struct RayHitInfo
{
	float		distToHit;
	float3		hitPos;
	float3		hitNormal;
	float2		hitTexcoord;
	int			matID;
};


__device__ int SphereRayCollisionCheck(const CudaSphere& sphere, const Ray& ray, RayHitInfo& hitInfo, float t)
{
	float3 dist = float3_sub(sphere.center, ray.origin);
	float B = dot(ray.dir, dist);
	float D = B * B - dot(dist, dist) + sphere.radius * sphere.radius;
	if (D < 0.0f)
		return FALSE;

	const float sqrtD = sqrtf(D);
	float t0 = B - sqrtD;
	float t1 = B + sqrtD;

	int retvalue = FALSE;
	if ((t0 > 0.1f) && (t0 < t)) 
	{
		t = t0;
		retvalue = TRUE;
	}
	if ((t1 > 0.1f) && (t1 < t)) 
	{
		t = t1;
		retvalue = TRUE;
	}

	if (retvalue)
	{
		hitInfo.distToHit = t;
		hitInfo.hitPos = float3_add(ray.origin, float3_mul(ray.dir, t));
		hitInfo.hitNormal = normalize(float3_sub(hitInfo.hitPos, sphere.center));
		hitInfo.hitTexcoord = make_float2(0.0f, 0.0f);
	}

	return retvalue;
}

__device__ float2 BoxCalcTexCoord(const CudaBox& box, float3 pos, int normAxis)
{
	float posU, posV;
	float2 minU, maxU;
	if (0 == normAxis)
	{
		minU.x = box.Minimum.z;
		minU.y = box.Minimum.y;
		maxU.x = box.Maximum.z;
		maxU.y = box.Maximum.y;

		posU = pos.z - minU.x;
		posV = pos.y - minU.y;
	}
	else if (1 == normAxis)
	{
		minU.x = box.Minimum.x;
		minU.y = box.Minimum.z;
		maxU.x = box.Maximum.x;
		maxU.y = box.Maximum.z;

		posU = pos.x - minU.x;
		posV = pos.z - minU.y;
	}
	else if (2 == normAxis)
	{
		minU.x = box.Minimum.x;
		minU.y = box.Minimum.y;
		maxU.x = box.Maximum.x;
		maxU.y = box.Maximum.y;

		posU = pos.x - minU.x;
		posV = pos.y - minU.y;
	}

	float2 diff = float2_sub(maxU, minU);
	return make_float2(posU / diff.x, 1.0f - (posV / diff.y));
}
__device__ int BoxRayCollisionCheck(const CudaBox& box, const Ray& ray, RayHitInfo& hitInfo, float t)
{
	float tmin, tmax, tymin, tymax, tzmin, tzmax;
	int normalAxis = 0,
		normalAxisMin = 0,
		normalAxisMax = 0;

	float3 MinMax[2];
	MinMax[0] = box.Minimum;
	MinMax[1] = box.Maximum;

	tmin  = (MinMax[    ray.posneg[0]].x - ray.origin.x) * ray.dir_inv.x;
	tmax  = (MinMax[1 - ray.posneg[0]].x - ray.origin.x) * ray.dir_inv.x;
	tymin = (MinMax[    ray.posneg[1]].y - ray.origin.y) * ray.dir_inv.y;
	tymax = (MinMax[1 - ray.posneg[1]].y - ray.origin.y) * ray.dir_inv.y;

	if ((tmin > tymax) || (tymin > tmax))
		return FALSE;

	if (tymin > tmin)
	{
		tmin = tymin;
		normalAxisMin = 1;
	}
	if (tymax < tmax)
	{
		tmax = tymax;
		normalAxisMax = 1;
	}

	tzmin = (MinMax[    ray.posneg[2]].z - ray.origin.z) * ray.dir_inv.z;
	tzmax = (MinMax[1 - ray.posneg[2]].z - ray.origin.z) * ray.dir_inv.z;

	if ((tmin > tzmax) || (tzmin > tmax))
		return FALSE;

	if (tzmin > tmin)
	{
		tmin = tzmin;
		normalAxisMin = 2;
	}
	if (tzmax < tmax)
	{
		tmax = tzmax;
		normalAxisMax = 2;
	}

	int result = FALSE;
	if ((tmin > 0.1f) && (tmin < t))
	{
		t = tmin;
		result = TRUE;
		normalAxis = normalAxisMin;
	}
	else if ((tmax > 0.1f) && (tmax < t))
	{
		t = tmax;
		result = TRUE;
		normalAxis = normalAxisMax;
	}

	if (result)
	{
		hitInfo.distToHit = t;
		hitInfo.hitPos = float3_add(ray.origin, float3_mul(ray.dir, t));

		float3 dir = float3_sub(hitInfo.hitPos, box.position);
		if (0 == normalAxis)
			hitInfo.hitNormal = make_float3(signf(dir.x), 0.0f, 0.0f);
		else if (1 == normalAxis)
			hitInfo.hitNormal = make_float3(0.0f, signf(dir.y), 0.0f);
		else if (2 == normalAxis)
			hitInfo.hitNormal = make_float3(0.0f, 0.0f, signf(dir.z));

		hitInfo.hitTexcoord = BoxCalcTexCoord(box, hitInfo.hitPos, normalAxis);
	}

	return result;
}

__device__ int PlaneRayCollisionCheck(const CudaPlane& plane, const Ray& ray, RayHitInfo& hitInfo, float t)
{
	int result = FALSE;

	float k0 = dot(plane.normal, ray.dir);
	if (k0 < 0.0f)
	{
		float k1 = dot(make_float4(plane.normal, plane.distance), make_float4(ray.origin, 1.0f));
		float dist = -k1 / k0;
		if (k1 > 0.0f && dist < t)
		{
			result = TRUE;
			hitInfo.distToHit = dist;
			hitInfo.hitPos = float3_add(ray.origin, float3_mul(ray.dir, dist));
			hitInfo.hitNormal = plane.normal;
			hitInfo.hitTexcoord = make_float2(dot(hitInfo.hitPos, plane.uAxis) * 0.0075625f,
											  dot(hitInfo.hitPos, plane.vAxis) * 0.0075625f);
		}
	}

	return result;
}


#endif	// __cuMath_cuh__
